#ifndef _BVH_ACCELERATOR_OCL_H_
#define _BVH_ACCELERATOR_OCL_H_

#include "primitive.h"
#include "triangle_mesh_ocl.h"

typedef struct 
#ifdef CL_KERNEL
	__attribute__ ((aligned (16)))
#endif
{
	bbox_t bounds;

	unsigned primitive_offset; // leaf
	unsigned second_child_offset;// interior

	unsigned char n_primitives;
	unsigned char axis;
	unsigned char pad_[6];
}
linear_bvh_node_t;

typedef struct{
	unsigned int linear_node_num;
	GLOBAL linear_bvh_node_t *nodes;
}ocl_bvh_accelerator_t;

INLINE void as_bvh_accelerator(GLOBAL const float* accelerator_data,ocl_bvh_accelerator_t* bvh)
{
	bvh->linear_node_num = as_uint(accelerator_data[0]);
	bvh->nodes = (GLOBAL linear_bvh_node_t*)(accelerator_data+4);
}

//intersection
//#define CL_KERNEL
INLINE bool bbox_intersectP(GLOBAL const bbox_t* bounds, const ray_t *ray, const vector3f_t *inv_dir)
{
	vector3f_t l1,l2;
	vsub(l1,bounds->pmin,ray->o);
	vmul(l1,l1,(*inv_dir));

	vsub(l2,bounds->pmax,ray->o);
	vmul(l2,l2,(*inv_dir));
	const vector3f_t t_near = fmin(l1,l2);
	const vector3f_t t_far = fmax(l1,l2);
	float t0 = max(max(max(t_near.x,t_near.y),max(t_near.x,t_near.z)),ray->mint);
	float t1 = min(min(min(t_far.x,t_far.y),min(t_far.x,t_far.z)),ray->maxt);

	return (t1 >= t0);
}

INLINE bool bbox_intersectP2(GLOBAL const bbox_t* bounds, const ray_t *ray, const vector3f_t *inv_dir, const unsigned *dir_is_neg)
{
	float tmin = ((dir_is_neg[0] == 0 ?bounds->pmin.x:bounds->pmax.x) - ray->o.x) * ((*inv_dir).x);
	float tmax = ((dir_is_neg[0] == 1 ?bounds->pmin.x:bounds->pmax.x) - ray->o.x) * (*inv_dir).x;

	float tymin = ((dir_is_neg[1] == 0 ?bounds->pmin.y:bounds->pmax.y) - ray->o.y) * (*inv_dir).y;
	float tymax = ((dir_is_neg[1] == 1 ?bounds->pmin.y:bounds->pmax.y) - ray->o.y) * (*inv_dir).y;
	if((tmin > tymax) || (tymin > tmax))
		return false;
	if (tymin > tmin ) tmin = tymin;
	if (tymax < tmax ) tmax = tymax;

	float tzmin = ((dir_is_neg[2] == 0 ?bounds->pmin.z:bounds->pmax.z) - ray->o.z) * (*inv_dir).z;
	float tzmax = ((dir_is_neg[2] == 1 ?bounds->pmin.z:bounds->pmax.z) - ray->o.z) * (*inv_dir).z;
	if((tmin > tzmax) || (tzmin > tmax))
		return false;
	if (tzmin > tmin)
		tmin = tzmin;
	if (tzmax < tmax)
		tmax = tzmax;
	return (tmin < ray->maxt) && (tmax > ray->mint);
}
#define MAX_BVH_DEPTH 32

INLINE bool bvh_intersect(const ocl_bvh_accelerator_t* bvh, ray_t *ray,GLOBAL const primitive_info_t* primitive_array, GLOBAL const float* shape_data, 
	GLOBAL intersection_t *isect,LOCAL int* local_stack)
{
	bool hit = false;
	vector3f_t inv_dir;
	vinit(inv_dir, 1.f/ray->d.x, 1.f/ ray->d.y, 1.f/ray->d.z);
	const unsigned dir_is_neg[3] = {inv_dir.x < 0, inv_dir.y < 0, inv_dir.z < 0};
	unsigned todo_offset = 0, node_num = 0;
#ifdef CL_KERNEL
	LOCAL int *todo = &local_stack[MAX_BVH_DEPTH * get_local_id(0)];
#else
	unsigned todo[MAX_BVH_DEPTH];
#endif
	//isect data
	float b1,b2;
	float thit;
	unsigned primitive_idx = 0xffffffff;
	triangle_t triangle;
	GLOBAL const float* intersected_triangle_memory_start;

	while(true)
	{
		const GLOBAL linear_bvh_node_t *node = bvh->nodes + node_num;
		if(bbox_intersectP(&node->bounds, ray, &inv_dir))
		{
			if (node->n_primitives > 0)
			{
				for(unsigned i = 0;i < node->n_primitives; ++i)
				{
					float btmp1,btmp2;
					triangle_t load_triangle;
					GLOBAL const float* triangle_mesh_start = load_triangle_points(shape_data,primitive_array[node->primitive_offset+i].shape_info,&load_triangle);
					if(intersect_triangle_simple(triangle_mesh_start,primitive_array[node->primitive_offset+i].shape_info.shape_type,
						&load_triangle,ray,
						&thit,&btmp1,&btmp2))
					{
						hit = true;
						ray->maxt = thit;
						primitive_idx = node->primitive_offset+i;
						intersected_triangle_memory_start = triangle_mesh_start;
						b1 = btmp1;
						b2 = btmp2;
						triangle = load_triangle;
					}
				}
				if(todo_offset == 0) break;
				node_num = todo[--todo_offset];
			}
			else
			{
				if(dir_is_neg[node->axis])
				{
					todo[todo_offset ++ ] = node_num + 1;
					node_num = node->second_child_offset;
				}
				else
				{
					todo[todo_offset ++ ] = node->second_child_offset;
					node_num = node_num + 1;
				}
			}
		}
		else
		{
			if(todo_offset == 0) break;
			node_num = todo[--todo_offset];
		}
	}
	if(hit)
	{
		point3f_t p;
		vclr(p);
		rpos(p,(*ray),(thit));
		
		differential_geometry_t dg;
		get_triangle_shading_info(intersected_triangle_memory_start, &triangle, primitive_array[primitive_idx].shape_info.shape_type, b1,b2,p, &dg);
		isect->ray_epsilon = 1e-3f * thit;
		isect->dg = dg;
	}
	isect->primitive_idx = primitive_idx;
	return hit;
}

INLINE bool bvh_intersectP(const ocl_bvh_accelerator_t* bvh, ray_t *ray,GLOBAL const primitive_info_t* primitive_array,GLOBAL const float* shape_data)
{
	if(bvh->linear_node_num == 0) return false;
	vector3f_t inv_dir;
	vinit(inv_dir, 1.f/ray->d.x, 1.f/ ray->d.y, 1.f/ray->d.z);
	unsigned dir_is_neg[3] = {inv_dir.x < 0, inv_dir.y < 0, inv_dir.z < 0};
	unsigned todo_offset = 0, node_num = 0;
	unsigned todo[MAX_BVH_DEPTH];
	triangle_t triangle;

	while(true)
	{
		const GLOBAL linear_bvh_node_t *node = bvh->nodes + node_num;
		if(bbox_intersectP2(&node->bounds, ray, &inv_dir,dir_is_neg))
		{
			if (node->n_primitives > 0)
			{
				for(unsigned i = 0;i < node->n_primitives; ++i)
				{
					load_triangle_points(shape_data,primitive_array[node->primitive_offset+i].shape_info,&triangle);
					
					if(intersect_triangleP(&triangle,ray))
					{
						return true;
					}
				}
				if(todo_offset == 0) break;
				node_num = todo[--todo_offset];
			}
			else
			{
				if(dir_is_neg[node->axis])
				{
					todo[todo_offset ++ ] = node_num + 1;
					node_num = node->second_child_offset;
				}
				else
				{
					todo[todo_offset ++ ] = node->second_child_offset;
					node_num = node_num + 1;
				}
			}
		}
		else
		{
			if(todo_offset == 0) break;
			node_num = todo[--todo_offset];
		}
	}
	return false;
}
#endif
